#define ROWMAJOR(x,y,w,h) y*w+x
#define COLMAJOR(x,y,w,h) x*h+y

struct pglVector
{
	float4 vec;
};

struct pglNormalVertex
{
	struct pglVector normal;
	struct pglVector position;
};

struct pglMatrix
{
	struct pglVector r[4];
};

struct pglMatrix matrixTranspose(
	struct pglMatrix matrix)
{
	struct pglMatrix out;

	out.r[0].vec = (float4)(matrix.r[0].vec.x,matrix.r[1].vec.x,matrix.r[2].vec.x,matrix.r[3].vec.x);
	out.r[1].vec = (float4)(matrix.r[0].vec.y,matrix.r[1].vec.y,matrix.r[2].vec.y,matrix.r[3].vec.y);
	out.r[2].vec = (float4)(matrix.r[0].vec.z,matrix.r[1].vec.z,matrix.r[2].vec.z,matrix.r[3].vec.z);
	out.r[3].vec = (float4)(matrix.r[0].vec.w,matrix.r[1].vec.w,matrix.r[2].vec.w,matrix.r[3].vec.w);

	return out;
}

struct pglVector vecMatMul(
	struct pglVector vector,
	struct pglMatrix matrix)
{
	struct pglVector out;

	struct pglMatrix transMatrix = matrixTranspose(matrix);

	out.vec = (float4)(
		dot(vector.vec,transMatrix.r[0].vec),
		dot(vector.vec,transMatrix.r[1].vec),
		dot(vector.vec,transMatrix.r[2].vec),
		dot(vector.vec,transMatrix.r[3].vec));

	return out;
}

// Per-vertex transform
__kernel 
void vertexTransform(
	__global struct pglNormalVertex *vertices,
	__global struct pglNormalVertex *outVertices,
	__global struct pglMatrix *matrices,
	uint numMatrices,
	uint4 viewport,
	__global uchar4 *pixels)
{
	uint gid = get_global_id(0);

	struct pglNormalVertex vertex = vertices[gid];
	uint width = viewport.z;
	uint height = viewport.w;
	
	// Transform by all matrices
	for(int i=0; i<numMatrices; i++)
	{
		struct pglMatrix matrix = matrices[i];
		vertex.position = vecMatMul(vertex.position,matrix);
	}

	// Divide by w
	vertex.position.vec.xyz /= vertex.position.vec.w;

	// Viewport Transform
	// NOTE viewport = (x,y,width,height)
	float wOver2 = (float)width/2.0f;
	float hOver2 = (float)height/2.0f;
	vertex.position.vec.x = (wOver2)*vertex.position.vec.x + ((float)viewport.x+wOver2);
	vertex.position.vec.y = (hOver2)*vertex.position.vec.y + ((float)viewport.y+hOver2);

	outVertices[gid] = vertex;

	uint screenx = (uint)vertex.position.vec.x;
	uint screeny = (uint)vertex.position.vec.y;
	if(screenx < 0 || screenx >= width || screeny < 0 || screeny >= height)
		return;

	// Draw vertices to screen (for debugging)
	uint pixelIndex = ROWMAJOR(screenx,screeny,width,height);
	pixels[pixelIndex] = (uchar4)(255,0,0,255);
}

// Per-triangle rasterizer
// Technically, called per-vertex, but discarding 2/3 of threads
// VERTICES ARE TRANSFORMED TO XY SPACE
__kernel
void rasterizer(
	__global uchar4* pixels,
	uint width,
	uint height,
	__global struct pglNormalVertex *vertices)
{
	uint gid = get_global_id(0);

	// only run for first vertex (per-triangle)
	if(gid%3 != 0)
		return;

	struct pglNormalVertex v1 = vertices[gid];
	struct pglNormalVertex v2 = vertices[gid+1];
	struct pglNormalVertex v3 = vertices[gid+2];

	float x1 = v1.position.vec.x;
	float x2 = v2.position.vec.x;
	float x3 = v3.position.vec.x;

	float y1 = v1.position.vec.y;
	float y2 = v2.position.vec.y;
	float y3 = v3.position.vec.y;

	// Bounding rectangle
    int minx = (int)min(x1, min(x2, x3));
    int maxx = (int)max(x1, max(x2, x3));
    int miny = (int)min(y1, min(y2, y3));
    int maxy = (int)max(y1, max(y2, y3));

	uint pixelIndex;
	int t1,t2,t3,t123;
	for(int y = miny; y < maxy; y++)
    {
        for(int x = minx; x < maxx; x++)
        {
            // When all half-space functions positive, pixel is in triangle
			t1 = ((x1 - x2) * (y - y1) - (y1 - y2) * (x - x1)) >= 0;
            t2 = ((x2 - x3) * (y - y2) - (y2 - y3) * (x - x2)) >= 0;
            t3 = ((x3 - x1) * (y - y3) - (y3 - y1) * (x - x3)) >= 0;
			t123 = t1*t2*t3;

			if(t123)
            {
				pixelIndex = ROWMAJOR(x,y,width,height);
				pixels[pixelIndex] = (uchar4)(0,255,0,255);
            }
        }
    }
}